//swap_rb.cu
#include <opencv2/core/cuda_types.hpp>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
extern "C" 
__global__ void swap_rb_kernel(const cv::cuda::PtrStepSz<uchar3> src, cv::cuda::PtrStep<uchar3> dst)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if (x < src.cols && y < src.rows)
    {
        uchar3 v = src(y, x);
        dst(y, x) = make_uchar3(v.z, v.y, v.x);
    }
}

extern "C" void swap_rb_caller(const cv::cuda::PtrStepSz<uchar3>& src, cv::cuda::PtrStep<uchar3> dst, cudaStream_t stream)
{
    dim3 block(32, 8);
    dim3 grid((src.cols + block.x - 1) / block.x, (src.rows + block.y - 1) / block.y);
    swap_rb_kernel <<<grid, block, 0, stream >>> (src, dst);
    if (stream == 0)
        cudaDeviceSynchronize();
}